/*
 *  Copyright 2010 INFN - APE group
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

#include <cuqu/cuqu.h>
#ifdef CUPRINTF_INCLUDED
#include <utils/cuPrintf.cu>
#endif

#ifndef MIN
#define MIN(A,B) (((A)<(B))?(A):(B))
#endif

#define CUQU_PACKED __attribute__((__packed__))

//#pragma pack(push)
//#pragma pack(1)

typedef struct param {
    enum {
        tag_none,
        tag_int,
        tag_uint,
        tag_float,
        tag_double
    } tag;
    int pad;
    union {
        int           si;
        unsigned int  ui;
        long          sl;
        unsigned long ul;
        float         f;
        double        d;
    };
}  param_t;
//#pragma pack(pop)

#define INIT_PARAM { param_t::tag_none, }

param_t p = INIT_PARAM;

typedef struct task {
    enum { cmd_none, cmd_open, cmd_close, cmd_write, cmd_quit };
    int cmd;
    int id;
    enum { max_params = 2 };
    param_t params[max_params];
} __attribute__((__packed__)) task_t;

#define INIT_TASK { task_t::cmd_none, 0, { INIT_PARAM, INIT_PARAM } }

task_t t = INIT_TASK;

typedef cuqu::host::queue<task_t> queue_t;
typedef queue_t::event_t event_t;
typedef cuqu::cuda_queue_t cuda_queue_t;

//----------------------------------------------------------------------

__device__ event_t event; // = INIT_GPU_EVENT(INIT_TASK);

__global__ void dequeue_tasks(cuda_queue_t cuda_q) 
{
    const unsigned int tid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
    const unsigned int bid = blockIdx.x + gridDim.x*blockIdx.y;
    int retcode = 0;
    //event_t e;
    int doquit = 0;

#ifdef CUPRINTF_CU
    //cuPrintfRestrict(0, CUPRINTF_UNRESTRICTED);
    //cuPrintfRestrict(CUPRINTF_UNRESTRICTED, 0);
    //cuPrintf("num_events=%d\n", num_events);
#endif

    while(!doquit) {
        retcode = cuda_q->timed_fetch(&event, 200 /*cuqu::timeout_infinite*/);
        if(cuqu::SUCCESS == retcode) {
            // execute task
            task_t tsk = event.u.t;
            if(0 == bid && 0 == tid) cudaPrintf("[%d:%d] tsk.cmd=%d\n", bid, tid, tsk.cmd);
            switch(tsk.cmd) {
            case task_t::cmd_open:
                break;
            case task_t::cmd_close:
                break;
            case task_t::cmd_write:
                break;
            case task_t::cmd_quit:
                if(0 == bid && 0 == tid) cudaPrintf("got cmd QUIT!\n");
                doquit = 1;
                break;
            default:
                cudaPrintf("ERROR: unsupported cmd=%d\n", tsk.cmd);
                break;
            }
        } else {
            cudaPrintf("ERROR: unexpected error with retcode=%d\n", retcode);
            doquit = 1;
            break;
        }
    }
}

//----------------------------------------------------------------------

void push_tasks(queue_t *q, size_t max_n_tasks)
{
    task_t tasks[] = {
        { task_t::cmd_open,  0, { param_t::tag_int, 0 } },
        { task_t::cmd_write, 0, { param_t::tag_int, 0 } },
        { task_t::cmd_write, 0, { param_t::tag_int, 0 } },
        { task_t::cmd_quit,  0, { param_t::tag_int, 0 } }
    };

    for(int k=0; k<MIN(max_n_tasks, sizeof(tasks)/sizeof(tasks[0])); ++k) {
        printf("# pushing task %d\n", k);
        int ret = q->timed_push(&tasks[k], cuqu::timeout_infinite);
        if(ret != cuqu::SUCCESS) {
            printf("ERROR: retcode %d, aborting at %d-th push\n", ret, k);
            break;
        }
    }
}

//----------------------------------------------------------------------

void test(int n_threads, int n_blocks, int nqevents)
{
#ifdef CUPRINTF_CU
    cudaPrintfInit();
#endif
    dim3 dimBlock(n_threads, 1, 1);
    dim3 dimGrid(n_blocks, 1, 1);

    printf("# sizeof(param_t)=%d\n", sizeof(param_t));
    printf("# sizeof(task_t)=%d\n", sizeof(task_t));

    printf("# queue CTOR with %d slots\n", nqevents); fflush(stdout);
    queue_t queue(nqevents, n_threads);
    queue_t *q = &queue;
    cuda_queue_t cuda_q = queue.get_cuda_ref();

    printf("# pushing at most %d tasks\n", nqevents); fflush(stdout);
    push_tasks(q, nqevents);

    cuquCheckMsg("aieee");

    // spawn GPU kernel
    printf("# spawning dequeue_tasks() kernel, n_threads=%d n_blocks=%d\n", n_threads, n_blocks); fflush(stdout);
    dequeue_tasks<<<dimGrid, dimBlock>>>(cuda_q);

    cuquCheckMsg("aieee");

    printf("# calling cudaThreadSynchronize()\n");
    cuquSafeCall(cudaThreadSynchronize());

    cuquCheckMsg("aieee");

    cuquSafeCall(cudaThreadSynchronize());
    cuquCheckMsg("# dequeue() execution failed");
    printf("kernel ended, reading back results\n"); fflush(stdout);
    

    cuqu::detail::raw_queue_t *raw_q = q->get_raw_ptr();
    printf("# raw_queue: nr=%d nw=%d\n", raw_q->nr, raw_q->nw); fflush(stdout);

#ifdef CUPRINTF_CU
    cudaPrintfDisplay(stdout, true);
    cudaPrintfEnd();
#endif

    // gives errors in queue DTOR
    //cudaThreadExit();    

    printf("# destroying queue object\n"); fflush(stdout);
}

//----------------------------------------------------------------------

int main(int argc, char* argv[])
{
    cudaDeviceProp deviceProp;
    int idev;

    //
    int nqevents = 8;
    int n_pump_events = nqevents;
    int n_threads = 32;
    int n_blocks = 14; // default 14 = # SM on C2050

    // CUDA initialization
    printf("# CUDA init\n"); fflush(stdout);
    idev = 0;
    cuquSafeCall(cudaSetDevice(idev));
    cuquSafeCall(cudaGetDeviceProperties(&deviceProp, idev));
    if(!deviceProp.canMapHostMemory) {
        fprintf(stderr, "ERROR: Device %d cannot map host memory!\n", idev);
        exit(EXIT_FAILURE);
    }
    int num_sm = deviceProp.multiProcessorCount;
    cuquSafeCall(cudaSetDeviceFlags(cudaDeviceMapHost));

    printf("# default n_blocks=%d\n", num_sm); fflush(stdout);
    n_blocks = num_sm;

#ifdef CUQU_DEBUG
    printf("WARNING: debug builds are not reliable, they may deadlock!\n"); fflush(stdout);
#endif

    // arg scan
    int opt;
    while ((opt = getopt(argc, argv, "he:b:t:")) != -1) {
        switch (opt) {
        case 'b':
            n_blocks = atoi(optarg);
            break;
        case 't':
            n_threads = atoi(optarg);
            break;
        case 'e':
            n_pump_events = atoi(optarg);
            //printf("parsed param -e %d\n", n_pump_events);
            break;
        default: /* '?' */
            fprintf(stderr, "Usage: %s [-e num_events][-b n_blocks][-t n_threads]\n",
                    argv[0]);
            exit(EXIT_FAILURE);
        }
    }

    test(n_threads, n_blocks, nqevents);
    printf("# test ended\n"); fflush(stdout);
  
    return EXIT_SUCCESS;
}

/*
 * Local variables:
 *  mode: c++
 *  c-indent-level: 4
 *  c-basic-offset: 4
 *  tab-width: 4
 *  indent-tabs-mode: nil
 * End:
 */
